#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h> 
#include "kernel.cuh"
#include <math.h>
#include <cutil.h>

using namespace std;

__global__ void RevertColorKernel(unsigned char *c, unsigned int size)
{
	unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
    unsigned int i = y * blockDim.x * gridDim.x + x;

    if(i < size)
	{
		c[i] = 255 - c[i];
	}
}

__global__ void LighterKernel(unsigned char *c, unsigned int size)
{
	unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
    unsigned int i = y * blockDim.x * gridDim.x + x;

    if(i < size)
	{
		if(255-c[i] < 20) c[i] = 255;
		else c[i] += 20;
	}
}

__global__ void Contrast(unsigned char *c, unsigned int size)
{
	unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
    unsigned int i = y * blockDim.x * gridDim.x + x;

    if(i < size)
	{
		c[i] = (sin(((((float)c[i]*180/255-90)*3.14/180)))+1)*127; //contrast
		//c[i] = (sin(((((float)c[i]*180/255)*3.14/180)))+1)*127; //some funy stuff
		//c[i] = ((int)c[i] + 20)%255;
	}
}

__global__ void HistogramGray(unsigned char *image, unsigned int* histo, unsigned int size)
{
	unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
    unsigned int i = y * blockDim.x * gridDim.x + x;

    if(i < size)
	{
		int j = i*3;
		atomicAdd(&histo[(int)((float)image[j] * 0.2989 + (float)image[j+1] * 0.5870 + (float)image[j+2] * 0.1140)], 1);
	}
}

__global__ void HistogramColor(unsigned char *image, unsigned int* histo, unsigned int size, int color)
{
	unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
    unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
    unsigned int i = y * blockDim.x * gridDim.x + x;

    if(i < size)
	{
		int j = i*3;
		atomicAdd(&histo[(int)image[j+color]], 1);
	}
}


void cudaErrorHandle(cudaError_t cudaStatus, char* message)
{
	if (cudaStatus != cudaSuccess) {
		fprintf(stderr, message);
		exit(1);
	}
}

Kernel::Kernel()
{
	dev_data = 0;

	cudaError_t cudaStatus = cudaSetDevice(0);
	cudaErrorHandle(cudaStatus, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
}

void Kernel::reset()
{
	 // cudaDeviceReset must be called before exiting in order for profiling and // tracing tools such as Parallel Nsight and Visual Profiler to show complete traces.
	cudaErrorHandle(cudaDeviceReset(), "cudaDeviceReset failed!");
}

float Kernel::convert(Image* image, OperationType cudaOper)
{
	copyImageToMemory(image);

	dim3 blockSize = dim3(32,32);
	dim3 gridSize(((image->width * image->getBytesPerPixel()) + blockSize.x - 1)/ blockSize.x, (image->height + blockSize.y - 1) / blockSize.y, 1);

	unsigned int kernelTime;
	cutCreateTimer(&kernelTime);
	cutResetTimer(kernelTime);

	switch(cudaOper)
	{
		case REVERT_COLOR: RevertColorKernel<<<gridSize, blockSize>>>(dev_data, image->getSize()); break;
		case LIGHTER: LighterKernel<<<gridSize, blockSize>>>(dev_data, image->getSize()); break;
		case CONTRAST: Contrast<<<gridSize, blockSize>>>(dev_data, image->getSize()); break;
	}

	cudaThreadSynchronize();
	cutStopTimer(kernelTime);

	cudaError_t cudaStatus = cudaDeviceSynchronize();
	cudaErrorHandle(cudaStatus, "cudaDeviceSynchronize returned error code %d after launching Kernel!\n");

	cudaStatus = cudaMemcpy(image->data, dev_data, image->getSize() * sizeof(char), cudaMemcpyDeviceToHost);
	cudaErrorHandle(cudaStatus, "cudaMemcpy failed Device -> Host!");

	return cutGetTimerValue(kernelTime);
}

unsigned int* Kernel::retrive(Image* image, RetriveType retType)
{
	unsigned int * dev_histo;

	copyImageToMemory(image);

	cudaError_t cudaStatus = cudaMalloc((void**)&dev_histo, 256 * sizeof(int));
	cudaErrorHandle(cudaStatus, "cudaMalloc failed!");

	unsigned int* histo = new unsigned int[256];
	for(int i=0;i<256;i++) histo[i]=0;

	cudaStatus = cudaMemcpy(dev_histo, histo, 256 * sizeof(int),  cudaMemcpyHostToDevice);
	cudaErrorHandle(cudaStatus, "cudaMemcpy failed Device -> Host!");

	dim3 blockSize = dim3(32,32);
	dim3 gridSize = dim3(((image->width) + blockSize.x - 1)/ blockSize.x, (image->height + blockSize.y - 1) / blockSize.y);

	unsigned int kernelTime;
	cutCreateTimer(&kernelTime);
	cutResetTimer(kernelTime);

	switch(retType)
	{
		case HISTOGRAM_GRAY: HistogramGray<<<gridSize, blockSize>>>(dev_data, dev_histo, image->width*image->height); break;
		case HISTOGRAM_RED: HistogramColor<<<gridSize, blockSize>>>(dev_data, dev_histo, image->width*image->height, 0); break;
		case HISTOGRAM_GREEN: HistogramColor<<<gridSize, blockSize>>>(dev_data, dev_histo, image->width*image->height, 1); break;
		case HISTOGRAM_BLUE: HistogramColor<<<gridSize, blockSize>>>(dev_data, dev_histo, image->width*image->height, 2); break;
	}

	cudaThreadSynchronize();
	cutStopTimer(kernelTime);

	cudaStatus = cudaDeviceSynchronize();
	cudaErrorHandle(cudaStatus, "cudaDeviceSynchronize returned error code %d after launching Kernel!\n");

	

	cudaStatus = cudaMemcpy(histo, dev_histo, 256 * sizeof(int), cudaMemcpyDeviceToHost);
	cudaErrorHandle(cudaStatus, "cudaMemcpy failed Device -> Host!");

	return histo;
}

void Kernel::copyImageToMemory(Image* image)
{
	if(image == imagePtrCpy && dev_data != NULL) return;

	cudaError_t cudaStatus = cudaMalloc((void**)&dev_data, image->getSize() * sizeof(char));
	cudaErrorHandle(cudaStatus, "cudaMalloc failed!");

	cudaStatus = cudaMemcpy(dev_data, image->data, image->getSize() * sizeof(char), cudaMemcpyHostToDevice);
	cudaErrorHandle(cudaStatus, "cudaMemcpy failed Host -> Device!");

	imagePtrCpy = image;
}






